// phalanx unit_test.hpp.cu header file

// Copyright (c) 2011 - 2011 Kohei Takahashi (Flast).
// Distributed under the MIT license. for more detail see COPYING.

#ifndef IG_PHALANX_TEST_UNIT_TEST_HPP_CU_ONCE_
#define IG_PHALANX_TEST_UNIT_TEST_HPP_CU_ONCE_

#include <phalanx/config.hpp.cu>

#include <cstddef>

#include <boost/test/unit_test.hpp>

#include <boost/preprocessor/cat.hpp>
#include <boost/preprocessor/stringize.hpp>

#include <phalanx/sandbox/test/detail/scoped_ptr.hpp.cu>

#include <phalanx/ptx/exit.hpp.cu>

namespace phalanx
{

namespace test_detail
{

class _device_test_launcher_impl
{
    typedef void ( *_kernel_type )( signed int * );

    cudaStream_t _launch_stream;

    const ::std::pair< const char *, ::std::size_t > _filename;
    const ::std::size_t _line;

protected:
    template < ::std::size_t N >
    PHALANX_CONFIG_ATTRIBUTE( host )
    _device_test_launcher_impl( const char ( &filename )[ N ], ::std::size_t line )
      : _filename( filename, N ), _line( line )
    {
        if ( cudaStreamCreate( &_launch_stream ) != cudaSuccess )
        { _launch_stream = 0; }
    }

    PHALANX_CONFIG_ATTRIBUTE( host )
    ~_device_test_launcher_impl( void )
    {
        if ( _launch_stream )
        { cudaStreamDestroy( _launch_stream ); }
    }

    // XXX: This funcction may be not portable.
    //      For more detail, see Boost.Test.UTF Reference.
    PHALANX_CONFIG_ATTRIBUTE( host )
    void
    _phalanx_test_message( const char *message )
    {
        using namespace ::boost::unit_test;
        const const_string filename( _filename.first, _filename.second - 1 );
        unit_test_log
          << log::begin( filename, _line )
          << lazy_ostream::instance()
          << message;
    }

#ifndef BOOST_TEST_PROD
#define LO_PHALANX_TEST_PASSPOINT() \
  ::boost::unit_test::unit_test_log.set_checkpoint( \
    ::boost::unit_test::const_string( \
      _l._filename.first, _l._filename.second - 1 ), \
    _l._filename.second )
#else
#define LO_PHALANX_TEST_PASSPOINT()
#endif

#define LO_PHALANX_TEST_TOOL_IMPL( line_, _pred, _TL, _CT, _message ) \
  do \
  { \
      const ::boost::test_tools::const_string filename( \
        _l._filename.first, _l._filename.second - 1 ); \
      LO_PHALANX_TEST_PASSPOINT(); \
      ::boost::test_tools::tt_detail:: \
      BOOST_PP_IF( 2, check_impl, check_frwd ) ( \
        BOOST_JOIN( BOOST_TEST_TOOL_PASS_PRED, 2 )( ( _pred ), _ ), \
        ::boost::unit_test::lazy_ostream::instance() << _message, \
        filename, line_, \
        ::boost::test_tools::tt_detail::_TL, \
        ::boost::test_tools::tt_detail::_CT \
        BOOST_JOIN( BOOST_TEST_TOOL_PASS_ARGS, 2 )( _ ) ); \
  } while ( ::boost::test_tools::dummy_cond )

#define LO_PHALANX_TEST_WARN( line_, _pred, _message ) \
  LO_PHALANX_TEST_TOOL_IMPL( line_, _pred, WARN, CHECK_PRED, _message )
#define LO_PHALANX_TEST_CHECK( line_, _pred, _message ) \
  LO_PHALANX_TEST_TOOL_IMPL( line_, _pred, CHECK, CHECK_PRED, _message )
#define LO_PHALANX_TEST_REQUIRE( line_, _pred, _message ) \
  LO_PHALANX_TEST_TOOL_IMPL( line_, _pred, REQUIRE, CHECK_PRED, _message )

#define LO_PHALANX_MESS_WARN( line_, _pred, _message ) \
  LO_PHALANX_TEST_TOOL_IMPL( line_, _pred, WARN, CHECK_MSG, _message )
#define LO_PHALANX_MESS_CHECK( line_, _pred, _message ) \
  LO_PHALANX_TEST_TOOL_IMPL( line_, _pred, CHECK, CHECK_MSG, _message )
#define LO_PHALANX_MESS_REQUIRE( line_, _pred, _message ) \
  LO_PHALANX_TEST_TOOL_IMPL( line_, _pred, REQUIRE, CHECK_MSG, _message )

#define LO_PHALANX_TOOL( _tool, line_, _pred ) \
  _tool( line_, _pred, BOOST_PP_STRINGIZE( _pred ) )

public:
    PHALANX_CONFIG_ATTRIBUTE( host )
    static void
    exec_kernel( _device_test_launcher_impl &_l, _kernel_type _kernel )
    {
        // FIXME: How to use mapped memory?
//#define _PHALANX_SANDBOX_TEST_USE_MAPPED_PTR
#ifdef _PHALANX_SANDBOX_TEST_USE_MAPPED_PTR
        if ( cudaSetDeviceFlags( cudaDeviceMapHost ) != cudaSuccess )
        { BOOST_FAIL( cudaGetErrorString( cudaPeekAtLastError() ) ); }
#endif

        // FIXME: Temporary values.
        const dim3 blocks( 1 );
        const dim3 threads( 32 );

        using namespace test_detail;
        // FIXME: Needs error recovery.
        typedef scoped_ptr< signed int, host > point_ptr_t;
#ifdef _PHALANX_SANDBOX_TEST_USE_MAPPED_PTR
        point_ptr_t failed_point( threads.x, cudaHostAllocMapped );
#else
        point_ptr_t failed_point( threads.x, cudaHostAllocPortable );
#endif

        const cudaError_t host_alloc_status = cudaPeekAtLastError();
        if ( host_alloc_status != cudaSuccess )
        { BOOST_FAIL( cudaGetErrorString( host_alloc_status ) ); }

#ifdef _PHALANX_SANDBOX_TEST_USE_MAPPED_PTR
        memset( failed_point.get(), -1, failed_point.instance_size() );
        typedef scoped_ptr< signed int, mapped > d_point_ptr_t;
        d_point_ptr_t d_failed_point( failed_point );
#else
        typedef scoped_ptr< signed int, device > d_point_ptr_t;
        d_point_ptr_t d_failed_point( threads.x );
        cudaMemset( d_failed_point.get(), -1, d_failed_point.instance_size() );
#endif

        _kernel<<< blocks, threads, 0, _l._launch_stream >>>( d_failed_point.get() );
        cudaStreamSynchronize( _l._launch_stream );

        const cudaError_t launch_status = cudaPeekAtLastError();
        if ( launch_status == cudaSuccess )
        {
            cudaMemcpy( failed_point.get(), d_failed_point.get(),
              d_failed_point.instance_size(), cudaMemcpyDeviceToHost );
        }
        else
        { BOOST_FAIL( cudaGetErrorString( launch_status ) ); }

        for ( int x = 0; x < threads.x; ++x )
        {
            // FIXME: Temporary message.
            LO_PHALANX_TEST_CHECK(
              failed_point[ x ],
              failed_point[ x ] == -1,
              "kernel execute" );
        }
    }

#undef LO_PHALANX_TOOL
#undef LO_PHALANX_MESS_REQUIRE
#undef LO_PHALANX_MESS_CHECK
#undef LO_PHALANX_MESS_WARN
#undef LO_PHALANX_TEST_REQUIRE
#undef LO_PHALANX_TEST_CHECK
#undef LO_PHALANX_TEST_WARN
#undef LO_PHALANX_TEST_TOOL_IMPL
#undef LO_PHALANX_TEST_PASSPOINT

};

#define LO_PHALANX_AUTO_TEST_CASE_ID( _name ) \
  BOOST_PP_CAT( \
    _phalanx_unnamed_test_launcher_ ## _name ## _, \
    BOOST_PP_CAT( __LINE__, _test_kernel ) )

#define LO_PHALANX_TEST_LAUNCHER_ID( _name ) \
  BOOST_PP_CAT( _device_test_launcher_impl_ ## _name ## _, __LINE__ )

#define PHALANX_AUTO_TEST_CASE( _name  ) \
  PHALANX_CONFIG_ATTRIBUTE( global ) \
  static void \
  LO_PHALANX_AUTO_TEST_CASE_ID( _name )( signed int * ); \
  \
  BOOST_AUTO_TEST_CASE( _name ) \
  { \
      using ::phalanx::test_detail::_device_test_launcher_impl; \
      class LO_PHALANX_TEST_LAUNCHER_ID( _name ) \
        : public _device_test_launcher_impl \
      { \
      public: \
          PHALANX_CONFIG_ATTRIBUTE( host ) \
          LO_PHALANX_TEST_LAUNCHER_ID( _name )( void ) \
            : _device_test_launcher_impl( __FILE__, __LINE__ ) {} \
      } t; \
      _device_test_launcher_impl::exec_kernel( t, LO_PHALANX_AUTO_TEST_CASE_ID( _name ) ); \
  } \
  PHALANX_CONFIG_ATTRIBUTE( global ) \
  static void \
  LO_PHALANX_AUTO_TEST_CASE_ID( _name )( \
    signed int *_phalanx_kerneltest_failed_point )

// FIXME: Temporary checker.
#define PHALANX_FAIL( mes ) \
  do \
  { \
    _phalanx_kerneltest_failed_point[ threadIdx.x ] = __LINE__; \
    ::phalanx::ptx::_exit(); \
  } while ( false )

// FIXME: Temporary checker.
#define PHALANX_REQUIRE( pred ) \
  do if ( !static_cast< bool >( pred ) ) \
  { \
      PHALANX_FAIL( # pred ); \
  } while ( false )

} // namespace test_detail

} // namespace phalanx

#endif // IG_PHALANX_TEST_UNIT_TEST_HPP_CU_ONCE_

